#include "kernel_operator.h"

using namespace AscendC;

#define ll long long

constexpr int32_t BUFFER_NUM = 1;

// constexpr uint64_t pre[65] = {0, 1ull,3ull,5ull,11ull,17ull,39ull,65ull,139ull,261ull,531ull,1025ull,2095ull,4097ull,8259ull,16405ull,32907ull,65537ull,131367ull,262145ull,524827ull,1048645ull,2098179ull,4194305ull,8390831ull,16777233ull,33558531ull,67109125ull,134225995ull,268435457ull,536887863ull,1073741825ull,2147516555ull,4294968325ull,8590000131ull,17179869265ull,34359871791ull,68719476737ull,137439215619ull,274877911045ull,549756338843ull,1099511627777ull,2199024312423ull,4398046511105ull,8796095120395ull,17592186061077ull,35184376283139ull,70368744177665ull,140737496778927ull,281474976710721ull,562949970199059ull,1125899906908165ull,2251799847243787ull,4503599627370497ull,9007199321981223ull,18014398509483025ull,36028797153190091ull,72057594038190085ull,144115188344291331ull,288230376151711745ull,576460752840837695ull,1152921504606846977ull,2305843010287435779ull,4611686018428436805ull,9223372039002292363ull};

template<typename T> class Lcm {
private:
	TPipe* pipe;
	TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX1;
	TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX2;
	TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueY;
	GlobalTensor<T> x1Gm;
	GlobalTensor<T> x2Gm;
	GlobalTensor<T> yGm;
	int blockIdx;
	ll offset;
	ll offset1;
	ll offset2;
	int indices[5];
	ll totalLength1;
	ll totalLength2;
	ll totalLength;
	int calcLength;
	int dimNum1;
	int dimNum2;
	int dimNum;
	int* dim1;
	int* dim2;
	int* dim;

public:
	__aicore__ inline Lcm() {}
	__aicore__ inline void Init(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, GM_ADDR workspace, TPipe* pipeIn, int totalLength1, int totalLength2, int totalLength, int dimNum1, int dimNum2, int dimNum, int* dim1, int* dim2, int* dim)
	{
		pipe = pipeIn;
		this->totalLength1 = totalLength1;
		this->totalLength2 = totalLength2;
		this->totalLength = totalLength;
		this->dimNum1 = dimNum1;
		this->dimNum2 = dimNum2;
		this->dimNum = dimNum;
		this->dim1 = dim1;
		this->dim2 = dim2;
		this->dim = dim;
		// 打印每个dimnum和dim
		// printf("dimNum1: %d, dimNum2: %d, dimNum: %d\n", dimNum1, dimNum2, dimNum);
		// for (int i = 0; i < dimNum; i++) {
		//     printf("dim[%d]: %d, dim1[%d]: %d, dim2[%d]: %d\n", i, dim[i], i, dim1[i], i, dim2[i]);
		// }
		
		blockIdx = GetBlockIdx();
		calcLength = dim[dimNum - 1];
		offset1 = blockIdx * calcLength;
		// 计算indices（每一维的索引）
		indices[dimNum - 1] = 0;
		int tmp = blockIdx;
		for (int i = dimNum - 2; i >= 0; i--) {
			indices[i] = tmp % dim[i];
			tmp /= dim[i];
		}
		// printf("blockIdx: %d, indices: ", blockIdx);
		// for (int i = 0; i < dimNum; i++) {
		//     printf("%d ", indices[i]);
		// }
		// printf("\n");

		// 计算x2索引
		offset2 = 0;
		ll stride2 = 1;
		for (int i = dimNum - 1; i >= 0; i--) {
			int idx = indices[i] % dim2[i]; // 处理广播
			offset2 += idx * stride2;
			stride2 *= dim2[i];
		}
		// printf("offset1: %d, offset2: %d\n", offset1, offset2);
		

		x1Gm.SetGlobalBuffer((__gm__ T *)x1 + offset1, calcLength);
		x2Gm.SetGlobalBuffer((__gm__ T *)x2 + offset2, dim2[dimNum - 1]);
		yGm.SetGlobalBuffer((__gm__ T *)y + offset1, calcLength);

		pipe->InitBuffer(inQueueX1, BUFFER_NUM, calcLength * sizeof(T));
		pipe->InitBuffer(inQueueX2, BUFFER_NUM, calcLength * sizeof(T));
		pipe->InitBuffer(outQueueY, BUFFER_NUM, calcLength * sizeof(T));
	}
	__aicore__ inline void Process()
	{
		// 如果只广播第二维
		// if (dimNum == 3 && dim2[1] == 1 && totalLength == totalLength2 * dim[1]) {
		// 	return;
		// }
		// for (int i = 0; i < dim[dimNum - 1]; i++) {
		//     // 获取x1和x2值并计算
		//     ll x1 = x1Gm(i);
		//     ll x2 = dim2[dimNum - 1] == 1 ? x2Gm(0) : x2Gm(i);
		//     ll y = get_lcm(x1, x2);
		//     yGm(i) = y;
		// }
		// return;

		DataCopyPadExtParams<T> padParams{false, 0, 0, 0};
		// VECIN和VECOUT是32B单位，其它都是1B单位
		DataCopyExtParams copyParams = {
			(uint16_t)1, 
			(uint32_t)(calcLength * sizeof(T)), 
			0,
			0,
			0
		};
		DataCopyExtParams copyParams2 = {
			(uint16_t)1, 
			(uint32_t)(1 * sizeof(T)), 
			0,
			0,
			0
		};
		if (dim2[dimNum - 1] == 1) {
			// 处理x2为1的情况
			LocalTensor<T> x1Local = inQueueX1.AllocTensor<T>();
			LocalTensor<T> x2Local = inQueueX2.AllocTensor<T>();
			DataCopyPad(x1Local, x1Gm, copyParams, padParams);
			DataCopyPad(x2Local, x2Gm, copyParams2, padParams);
			inQueueX1.EnQue(x1Local);
			inQueueX2.EnQue(x2Local);
			x1Local = inQueueX1.DeQue<T>();
			x2Local = inQueueX2.DeQue<T>();
			T x2 = x2Local(0);
			LocalTensor<T> yLocal = outQueueY.AllocTensor<T>();

			for (int i = 0; i < calcLength; i++) {
				T x1 = x1Local(i);
				T y = get_lcm(x1, x2);
				yLocal(i) = y;
			}

			inQueueX1.FreeTensor(x1Local);
			inQueueX2.FreeTensor(x2Local);
			outQueueY.EnQue(yLocal);
			yLocal = outQueueY.DeQue<T>();
			DataCopyPad(yGm, yLocal, copyParams);
			outQueueY.FreeTensor(yLocal);
		} else {
			// 处理x2为多维的情况
			LocalTensor<T> x1Local = inQueueX1.AllocTensor<T>();
			LocalTensor<T> x2Local = inQueueX2.AllocTensor<T>();
			DataCopyPad(x1Local, x1Gm, copyParams, padParams);
			DataCopyPad(x2Local, x2Gm, copyParams, padParams);
			inQueueX1.EnQue(x1Local);
			inQueueX2.EnQue(x2Local);
			x1Local = inQueueX1.DeQue<T>();
			x2Local = inQueueX2.DeQue<T>();
			LocalTensor<T> yLocal = outQueueY.AllocTensor<T>();

			for (int i = 0; i < calcLength; i++) {
				T x1 = x1Local(i);
				T x2 = x2Local(i);
				T y = get_lcm(x1, x2);
				yLocal(i) = y;
			}

			inQueueX1.FreeTensor(x1Local);
			inQueueX2.FreeTensor(x2Local);
			outQueueY.EnQue(yLocal);
			yLocal = outQueueY.DeQue<T>();
			DataCopyPad(yGm, yLocal, copyParams);
			outQueueY.FreeTensor(yLocal);
		}
	}
	__aicore__ inline T abs(T x) {
		return x < 0 ? -x : x;
	}
	__aicore__ inline T gcd(T a, T b) {
		while (b != 0) {
			T tmp = a % b;
			a = b;
			b = tmp;
		}
		// return abs(a);
		return a;
	}
	__aicore__ inline T stein_gcd(T a, T b) {
		// if (a == 0) return b;
		// if (b == 0) return a;
		// a = abs(a);
		// b = abs(b);
		T shift = ScalarGetSFFValue<1>(a | b);
		a >>= ScalarGetSFFValue<1>(a);
		do {
			b >>= ScalarGetSFFValue<1>(b);
			// if(a <= 64 && b <= 64){
			// 	a = 64 - ScalarCountLeadingZero(pre[a] & pre[b]);
			// 	break;
			// }
			if (a > b) {
				// swap(a, b);
				a ^= b ^= a ^= b;
			}
			b -= a;
		} while (b);
		return a << shift;
	}
	__aicore__ inline T get_lcm(T a, T b) {
		if (a == 0 || b == 0) return 0; // 避免除以0
		a = abs(a);
		b = abs(b);
		// 如果是int64，则最后结果取绝对值
		if constexpr (std::is_same_v<T, int64_t>) {
			return abs((a / stein_gcd(a, b)) * b);
		} else {
			return (a / stein_gcd(a, b)) * b;
		}
	}
};


template<typename T> class Lcm_SP {
private:
	TPipe* pipe;
	TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX1;
	TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX2;
	TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueY;
    TBuf<TPosition::VECCALC> tBufMask, tBufFloat;
	GlobalTensor<T> x1Gm;
	GlobalTensor<T> x2Gm;
	GlobalTensor<T> yGm;
	int blockIdx;
	int blockNum;
	ll offset;
	ll offset1;
	ll offset2;
	int indices[5];
	ll totalLength1;
	ll totalLength2;
	ll totalLength;
	int tileLength;
	int dimNum1;
	int dimNum2;
	int dimNum;
	int* dim1;
	int* dim2;
	int* dim;

public:
	__aicore__ inline Lcm_SP() {}
	__aicore__ inline void Init(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, GM_ADDR workspace, TPipe* pipeIn, int totalLength1, int totalLength2, int totalLength, int dimNum1, int dimNum2, int dimNum, int* dim1, int* dim2, int* dim)
	{
		pipe = pipeIn;
		this->totalLength1 = totalLength1;
		this->totalLength2 = totalLength2;
		this->totalLength = totalLength;
		this->dimNum1 = dimNum1;
		this->dimNum2 = dimNum2;
		this->dimNum = dimNum;
		this->dim1 = dim1;
		this->dim2 = dim2;
		this->dim = dim;
		// 打印每个dimnum和dim
		// printf("dimNum1: %d, dimNum2: %d, dimNum: %d\n", dimNum1, dimNum2, dimNum);
		// for (int i = 0; i < dimNum; i++) {
		//     printf("dim[%d]: %d, dim1[%d]: %d, dim2[%d]: %d\n", i, dim[i], i, dim1[i], i, dim2[i]);
		// }
		
		blockIdx = GetBlockIdx();
		blockNum = GetBlockNum();

		tileLength = 4096;
		if (tileLength > totalLength / blockNum) {
			// totalLength / blockNum，向上256字节对齐
			tileLength = (totalLength / blockNum * sizeof(T) + 255) / 256 * 256 / sizeof(T);
		}

		x1Gm.SetGlobalBuffer((__gm__ T *)x1, totalLength);
		x2Gm.SetGlobalBuffer((__gm__ T *)x2, totalLength);
		yGm.SetGlobalBuffer((__gm__ T *)y, totalLength);

		pipe->InitBuffer(inQueueX1, BUFFER_NUM, tileLength * sizeof(T));
		pipe->InitBuffer(inQueueX2, BUFFER_NUM, tileLength * sizeof(T));
		pipe->InitBuffer(outQueueY, BUFFER_NUM, tileLength * sizeof(T));
        if constexpr(std::is_same_v<T, int32_t>) {
            pipe->InitBuffer(tBufMask, 1 * tileLength * sizeof(uint8_t));
            pipe->InitBuffer(tBufFloat, 4 * tileLength * sizeof(float));
        }
	}
	__aicore__ inline void Process()
	{
        // if constexpr(std::is_same_v<T, int16_t>) {
        // } else if constexpr(std::is_same_v<T, int32_t>) {
        // } else {
		// 	return;
		// }
        #define min(x, y) ((x) < (y) ? (x) : (y))
        for (size_t offset = blockIdx * tileLength; offset < totalLength; offset += blockNum * tileLength) {
            int calcLength = min(tileLength, totalLength - offset);
            Compute(offset, calcLength);
        }
	}
	__aicore__ inline void Compute(size_t offset, int calcLength)
	{
		DataCopyPadExtParams<T> padParams{false, 0, 0, 0};
		// VECIN和VECOUT是32B单位，其它都是1B单位
		DataCopyExtParams copyParams = {
			(uint16_t)1, 
			(uint32_t)(tileLength * sizeof(T)), 
			0,
			0,
			0
		};
		// 处理x2为多维的情况
		LocalTensor<T> x1Local = inQueueX1.AllocTensor<T>();
		LocalTensor<T> x2Local = inQueueX2.AllocTensor<T>();
		DataCopyPad(x1Local, x1Gm[offset], copyParams, padParams);
		DataCopyPad(x2Local, x2Gm[offset], copyParams, padParams);
		inQueueX1.EnQue(x1Local);
		inQueueX2.EnQue(x2Local);
		x1Local = inQueueX1.DeQue<T>();
		x2Local = inQueueX2.DeQue<T>();
		LocalTensor<T> yLocal = outQueueY.AllocTensor<T>();

		if constexpr(std::is_same_v<T, int32_t>) {
            LcmLikely32(yLocal, x1Local, x2Local);
        } else {
            for (int i = 0; i < calcLength; i++) {
				T x1 = x1Local(i);
				T x2 = x2Local(i);
				T y = get_lcm(x1, x2);
                yLocal(i) = y;
            }
        }

		inQueueX1.FreeTensor(x1Local);
		inQueueX2.FreeTensor(x2Local);
		outQueueY.EnQue(yLocal);
		yLocal = outQueueY.DeQue<T>();
		DataCopyPad(yGm[offset], yLocal, copyParams);
		outQueueY.FreeTensor(yLocal);
	}
	__aicore__ inline T abs(T x) {
		return x < 0 ? -x : x;
	}
	__aicore__ inline T gcd(T a, T b) {
		while (b != 0) {
			T tmp = b;
			b = a % b;
			a = tmp;
		}
		return a;
	}
	__aicore__ inline T stein_gcd(T a, T b) {
		// if (a == 0) return b;
		// if (b == 0) return a;
		// a = abs(a);
		// b = abs(b);
		T shift = ScalarGetSFFValue<1>(a | b);
		a >>= ScalarGetSFFValue<1>(a);
		do {
			b >>= ScalarGetSFFValue<1>(b);
			// if(a <= 64 && b <= 64){
			// 	a = 64 - ScalarCountLeadingZero(pre[a] & pre[b]);
			// 	break;
			// }
			if (a > b) {
				// swap(a, b);
				a ^= b ^= a ^= b;
			}
			b -= a;
		} while (b);
		return a << shift;
	}
	__aicore__ inline T get_lcm(T a, T b) {
		if (a == 0 || b == 0) return 0; // 避免除以0
		a = abs(a);
		b = abs(b);
		// 如果是int64，则最后结果取绝对值
		if constexpr (std::is_same_v<T, int64_t>) {
			return abs((a / stein_gcd(a, b)) * b);
		} else {
			return (a / stein_gcd(a, b)) * b;
		}
	}
    __aicore__ inline void LcmLikely32(const LocalTensor<int32_t>& c, const LocalTensor<int32_t>& a, const LocalTensor<int32_t>& b) {
        auto mask = tBufMask.Get<uint8_t>();
		auto f_a = tBufFloat.Get<float>();
		auto f_b = f_a[tileLength + 32];
		auto f_c = f_a[(tileLength + 32) * 2];
		auto q = f_a[(tileLength + 32) * 3];

		// 打印前5个的a, b
		// for (int i = 0; i < 5; i++) {
		// 	int a_v = a.GetValue(i);
		// 	int b_v = b.GetValue(i);
		// 	printf("Input A[%d]: %d, Input B[%d]: %d\n",
		// 		i, a_v, i, b_v);
		// }

		Cast(f_a, a, RoundMode::CAST_TRUNC, tileLength);
		Cast(f_b, b, RoundMode::CAST_TRUNC, tileLength);

		// 取绝对值
		Abs(f_a, f_a, tileLength);
		Abs(f_b, f_b, tileLength);

		Mul(f_c, f_a, f_b, tileLength);

		Min(q, f_a, f_b, tileLength);
		Max(f_b, f_a, f_b, tileLength);
		// DataCopy(q, f_a, tileLength);

		// 打印前5个的f_a, f_b, q
		// for (int i = 0; i < 5; i++) {
		// 	float f_a_v = f_a.GetValue(i);
		// 	float f_b_v = f_b.GetValue(i);
		// 	float f_q_v = q.GetValue(i);
		// 	printf("Before compute, Input F_A[%d]: %f, Input F_B[%d]: %f, Input F_Q[%d]: %f\n",
		// 		i, f_a_v, i, f_b_v, i, f_q_v);
		// }

		#define REP 11
        for (int i = 0; i < REP; i++) {
			if (i % 3 == 0) {
				// q小b大
				CompareScalar(mask, q, (float)0.0f, CMPMODE::EQ, tileLength);
				Select(q, mask, f_b, q, SELMODE::VSEL_TENSOR_TENSOR_MODE, tileLength);
				Fmod(f_a, f_b, q, tileLength);
			} else if (i % 3 == 1) {
				// a小q大
				CompareScalar(mask, f_a, (float)0.0f, CMPMODE::EQ, tileLength);
				Select(f_a, mask, q, f_a, SELMODE::VSEL_TENSOR_TENSOR_MODE, tileLength);
				Fmod(f_b, q, f_a, tileLength);
			} else {
				// b小a大
				CompareScalar(mask, f_b, (float)0.0f, CMPMODE::EQ, tileLength);
				Select(f_b, mask, f_a, f_b, SELMODE::VSEL_TENSOR_TENSOR_MODE, tileLength);
				Fmod(q, f_a, f_b, tileLength);
			}
		}
		if (REP % 3 == 0) {
			// q小b大
			Div(f_c, f_c, f_b, tileLength);
		} else if (REP % 3 == 1) {
			// a小q大
			Div(f_c, f_c, q, tileLength);
		} else {
			// b小a大
			Div(f_c, f_c, f_a, tileLength);
		}
		// 收尾：理论上 REP16 次后所有 b==0，无需标量回退，现在f_a是gcd结果
		// 打印前5个的f_a, f_b, q
		// for (int i = 0; i < 5; i++) {
		// 	float f_a_v = f_a.GetValue(i);
		// 	float f_b_v = f_b.GetValue(i);
		// 	float f_q_v = q.GetValue(i);
		// 	printf("Input F_A[%d]: %f, Input F_B[%d]: %f, Input F_Q[%d]: %f\n",
		// 		i, f_a_v, i, f_b_v, i, f_q_v);
		// }

		Cast(c, f_c, RoundMode::CAST_TRUNC, tileLength);

		// 打印前5个的c
		// for (int i = 0; i < 5; i++) {
		// 	int c_v = c.GetValue(i);
		// 	printf("C[%d]: %d\n", i, c_v);
		// }
    }
};


template<typename T> class Lcm_CS3 {
private:
	TPipe* pipe;
	TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX1;
	TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX2;
	TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueY;
    TBuf<TPosition::VECCALC> tBufMask, tBufFloat;
	GlobalTensor<T> x1Gm;
	GlobalTensor<T> x2Gm;
	GlobalTensor<T> yGm;
	int blockIdx;
	int blockNum;
	ll offset;
	ll offset1;
	ll offset2;
	int indices[5];
	ll totalLength1;
	ll totalLength2;
	ll totalLength;
	int tileLength;
	int calcLength;
	int dimNum1;
	int dimNum2;
	int dimNum;
	int* dim1;
	int* dim2;
	int* dim;

public:
	__aicore__ inline Lcm_CS3() {}
	__aicore__ inline void Init(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, GM_ADDR workspace, TPipe* pipeIn, int totalLength1, int totalLength2, int totalLength, int dimNum1, int dimNum2, int dimNum, int* dim1, int* dim2, int* dim)
	{
		pipe = pipeIn;
		this->totalLength1 = totalLength1;
		this->totalLength2 = totalLength2;
		this->totalLength = totalLength;
		this->dimNum1 = dimNum1;
		this->dimNum2 = dimNum2;
		this->dimNum = dimNum;
		this->dim1 = dim1;
		this->dim2 = dim2;
		this->dim = dim;
		// 打印每个dimnum和dim
		// printf("dimNum1: %d, dimNum2: %d, dimNum: %d\n", dimNum1, dimNum2, dimNum);
		// for (int i = 0; i < dimNum; i++) {
		//     printf("dim[%d]: %d, dim1[%d]: %d, dim2[%d]: %d\n", i, dim[i], i, dim1[i], i, dim2[i]);
		// }
		blockIdx = GetBlockIdx();
		calcLength = dim[dimNum - 1];
		// 上对齐到256字节
		tileLength = (calcLength * sizeof(T) + 255) / 256 * 256 / sizeof(T);
		offset1 = blockIdx * calcLength;
		// 计算indices（每一维的索引）
		indices[dimNum - 1] = 0;
		int tmp = blockIdx;
		for (int i = dimNum - 2; i >= 0; i--) {
			indices[i] = tmp % dim[i];
			tmp /= dim[i];
		}
		// printf("blockIdx: %d, indices: ", blockIdx);
		// for (int i = 0; i < dimNum; i++) {
		//     printf("%d ", indices[i]);
		// }
		// printf("\n");

		// 计算x2索引
		offset2 = 0;
		ll stride2 = 1;
		for (int i = dimNum - 1; i >= 0; i--) {
			int idx = indices[i] % dim2[i]; // 处理广播
			offset2 += idx * stride2;
			stride2 *= dim2[i];
		}
		// printf("offset1: %d, offset2: %d\n", offset1, offset2);
		

		x1Gm.SetGlobalBuffer((__gm__ T *)x1 + offset1, tileLength);
		x2Gm.SetGlobalBuffer((__gm__ T *)x2 + offset2, tileLength);
		yGm.SetGlobalBuffer((__gm__ T *)y + offset1, tileLength);

		pipe->InitBuffer(inQueueX1, BUFFER_NUM, tileLength * sizeof(T));
		pipe->InitBuffer(inQueueX2, BUFFER_NUM, tileLength * sizeof(T));
		pipe->InitBuffer(outQueueY, BUFFER_NUM, tileLength * sizeof(T));
		pipe->InitBuffer(tBufMask, 1 * tileLength * sizeof(uint8_t));
		pipe->InitBuffer(tBufFloat, 4 * tileLength * sizeof(float));
	}
	__aicore__ inline void Process()
	{
		DataCopyPadExtParams<T> padParams{false, 0, 0, 0};
		// VECIN和VECOUT是32B单位，其它都是1B单位
		DataCopyExtParams copyParams = {
			(uint16_t)1, 
			(uint32_t)(tileLength * sizeof(T)), 
			0,
			0,
			0
		};
		// 处理x2为多维的情况
		LocalTensor<T> x1Local = inQueueX1.AllocTensor<T>();
		LocalTensor<T> x2Local = inQueueX2.AllocTensor<T>();
		DataCopyPad(x1Local, x1Gm, copyParams, padParams);
		DataCopyPad(x2Local, x2Gm, copyParams, padParams);
		inQueueX1.EnQue(x1Local);
		inQueueX2.EnQue(x2Local);
		x1Local = inQueueX1.DeQue<T>();
		x2Local = inQueueX2.DeQue<T>();
		LocalTensor<T> yLocal = outQueueY.AllocTensor<T>();

		Lcm64(yLocal, x1Local, x2Local);

		inQueueX1.FreeTensor(x1Local);
		inQueueX2.FreeTensor(x2Local);
		outQueueY.EnQue(yLocal);
		yLocal = outQueueY.DeQue<T>();
		copyParams = {
			(uint16_t)1, 
			(uint32_t)(calcLength * sizeof(T)), 
			0,
			0,
			0
		};
		DataCopyPad(yGm, yLocal, copyParams);
		outQueueY.FreeTensor(yLocal);
	}
	__aicore__ inline T abs(T x) {
		return x < 0 ? -x : x;
	}
	__aicore__ inline T gcd(T a, T b) {
		while (b != 0) {
			T tmp = b;
			b = a % b;
			a = tmp;
		}
		return a;
	}
	__aicore__ inline T stein_gcd_to24(T a, T b, const LocalTensor<float>& f_a, const LocalTensor<float>& f_b, int index) {
		if (a == 0 || b == 0) return 0;
		T shift = ScalarGetSFFValue<1>(a | b);
		a >>= ScalarGetSFFValue<1>(a);
		do {
			b >>= ScalarGetSFFValue<1>(b);
			if (a > b) {
				// swap(a, b);
				a ^= b ^= a ^= b;
			}
			b -= a;
		} while (b >= (1<<23) || a >= (1<<23));
		// 要b小a大
		if (a < b) {
			// swap(a, b);
			a ^= b ^= a ^= b;
		}
		f_a(index) = (float)a;
		f_b(index) = (float)b;
		return shift;
	}
    __aicore__ inline void Lcm64(const LocalTensor<T>& cLocal, const LocalTensor<T>& aLocal, const LocalTensor<T>& bLocal) {
		auto f_a = tBufFloat.Get<float>();
		auto f_b = f_a[tileLength];
		auto f_c = f_a[tileLength * 2];
		auto q = f_a[tileLength * 3];
		T va[1536], vb[1536];

		// 打印前10个的a, b
		// if (blockIdx == 0) {
		// 	for (int i = 0; i < 10; i++) {
		// 		ll a_v = aLocal.GetValue(i);
		// 		ll b_v = bLocal.GetValue(i);
		// 		printf("Input A[%d]: %lld, Input B[%d]: %lld\n",
		// 			i, a_v, i, b_v);
		// 	}
		// }

		// for (int i = 0; i < tileLength; i++) {
		// 	T a = aLocal(i);
		// 	T b = bLocal(i);
		// 	a = a < 0 ? -a : a;
		// 	b = b < 0 ? -b : b;
		// 	T shift = stein_gcd_to24(a, b, f_a, f_b, i);
		// 	va[i] = a >> shift;
		// 	vb[i] = b;
		// 	// aLocal(i) = (a >> shift);
		// 	// bLocal(i) = b;
		// 	// b小a大
		// 	// f_a(i) = (float)a;
		// 	// f_b(i) = (float)b;
		// }
		{
			int i = 0;
			for (; i <= tileLength - 10; i += 10) {
				// 从aLocal取值
				T a00 = aLocal(i + 0);
				T a01 = aLocal(i + 1);
				T a02 = aLocal(i + 2);
				T a03 = aLocal(i + 3);
				T a04 = aLocal(i + 4);
				T a05 = aLocal(i + 5);
				T a06 = aLocal(i + 6);
				T a07 = aLocal(i + 7);
				T a08 = aLocal(i + 8);
				T a09 = aLocal(i + 9);

				// 从bLocal取值
				T b00 = bLocal(i + 0);
				T b01 = bLocal(i + 1);
				T b02 = bLocal(i + 2);
				T b03 = bLocal(i + 3);
				T b04 = bLocal(i + 4);
				T b05 = bLocal(i + 5);
				T b06 = bLocal(i + 6);
				T b07 = bLocal(i + 7);
				T b08 = bLocal(i + 8);
				T b09 = bLocal(i + 9);

				// 取绝对值
				a00 = a00 < 0 ? -a00 : a00;
				a01 = a01 < 0 ? -a01 : a01;
				a02 = a02 < 0 ? -a02 : a02;
				a03 = a03 < 0 ? -a03 : a03;
				a04 = a04 < 0 ? -a04 : a04;
				a05 = a05 < 0 ? -a05 : a05;
				a06 = a06 < 0 ? -a06 : a06;
				a07 = a07 < 0 ? -a07 : a07;
				a08 = a08 < 0 ? -a08 : a08;
				a09 = a09 < 0 ? -a09 : a09;

				b00 = b00 < 0 ? -b00 : b00;
				b01 = b01 < 0 ? -b01 : b01;
				b02 = b02 < 0 ? -b02 : b02;
				b03 = b03 < 0 ? -b03 : b03;
				b04 = b04 < 0 ? -b04 : b04;
				b05 = b05 < 0 ? -b05 : b05;
				b06 = b06 < 0 ? -b06 : b06;
				b07 = b07 < 0 ? -b07 : b07;
				b08 = b08 < 0 ? -b08 : b08;
				b09 = b09 < 0 ? -b09 : b09;

				// 计算
				T shift00 = stein_gcd_to24(a00, b00, f_a, f_b, i + 0);
				T shift01 = stein_gcd_to24(a01, b01, f_a, f_b, i + 1);
				T shift02 = stein_gcd_to24(a02, b02, f_a, f_b, i + 2);
				T shift03 = stein_gcd_to24(a03, b03, f_a, f_b, i + 3);
				T shift04 = stein_gcd_to24(a04, b04, f_a, f_b, i + 4);
				T shift05 = stein_gcd_to24(a05, b05, f_a, f_b, i + 5);
				T shift06 = stein_gcd_to24(a06, b06, f_a, f_b, i + 6);
				T shift07 = stein_gcd_to24(a07, b07, f_a, f_b, i + 7);
				T shift08 = stein_gcd_to24(a08, b08, f_a, f_b, i + 8);
				T shift09 = stein_gcd_to24(a09, b09, f_a, f_b, i + 9);

				// 更新va
				va[i + 0] = a00 >> shift00;
				va[i + 1] = a01 >> shift01;
				va[i + 2] = a02 >> shift02;
				va[i + 3] = a03 >> shift03;
				va[i + 4] = a04 >> shift04;
				va[i + 5] = a05 >> shift05;
				va[i + 6] = a06 >> shift06;
				va[i + 7] = a07 >> shift07;
				va[i + 8] = a08 >> shift08;
				va[i + 9] = a09 >> shift09;

				// 更新vb
				vb[i + 0] = b00;
				vb[i + 1] = b01;
				vb[i + 2] = b02;
				vb[i + 3] = b03;
				vb[i + 4] = b04;
				vb[i + 5] = b05;
				vb[i + 6] = b06;
				vb[i + 7] = b07;
				vb[i + 8] = b08;
				vb[i + 9] = b09;
			}
			while (i < tileLength) {
				T a = aLocal(i);
				T b = bLocal(i);
				a = a < 0 ? -a : a;
				b = b < 0 ? -b : b;
				T shift = stein_gcd_to24(a, b, f_a, f_b, i);
				va[i] = a >> shift;
				vb[i] = b;
				i++;
			}
		}

		// 打印前10个的a, b, fa, fb
		// if (blockIdx == 0) {
		// 	for (int i = 0; i < 10; i++) {
		// 		ll a_v = aLocal.GetValue(i);
		// 		ll b_v = bLocal.GetValue(i);
		// 		float fa_v = f_a.GetValue(i);
		// 		float fb_v = f_b.GetValue(i);
		// 		printf("Input A[%d]: %lld, Input B[%d]: %lld, FA[%d]: %f, FB[%d]: %f\n",
		// 			i, a_v, i, b_v, i, fa_v, i, fb_v);
		// 	}
		// }

        auto mask = tBufMask.Get<uint8_t>();
		#define REP64 25
        for (int i = 0; i < REP64; i++) {
			if (i % 3 == 0) {
				// b小a大
				CompareScalar(mask, f_b, (float)0.0f, CMPMODE::EQ, tileLength);
				Select(f_b, mask, f_a, f_b, SELMODE::VSEL_TENSOR_TENSOR_MODE, tileLength);
				Fmod(q, f_a, f_b, tileLength);
			} else if (i % 3 == 1) {
				// q小b大
				CompareScalar(mask, q, (float)0.0f, CMPMODE::EQ, tileLength);
				Select(q, mask, f_b, q, SELMODE::VSEL_TENSOR_TENSOR_MODE, tileLength);
				Fmod(f_a, f_b, q, tileLength);
			} else {
				// a小q大
				CompareScalar(mask, f_a, (float)0.0f, CMPMODE::EQ, tileLength);
				Select(f_a, mask, q, f_a, SELMODE::VSEL_TENSOR_TENSOR_MODE, tileLength);
				Fmod(f_b, q, f_a, tileLength);
			}
		}

		// 打印前10个的fa, fb, q, a, b
		// if (blockIdx == 0) {
		// 	for (int i = 0; i < 10; i++) {
		// 		float fa_v = f_a.GetValue(i);
		// 		float fb_v = f_b.GetValue(i);
		// 		float q_v = q.GetValue(i);
		// 		ll a_v = aLocal.GetValue(i);
		// 		ll b_v = bLocal.GetValue(i);
		// 		printf("FA[%d]: %f, FB[%d]: %f, Q[%d]: %f, A[%d]: %lld, B[%d]: %lld\n",
		// 			i, fa_v, i, fb_v, i, q_v, i, a_v, i, b_v);
		// 	}
		// }

		// for (int i = 0; i < tileLength; i++) {
		// 	if constexpr (REP64 % 3 == 0) {
		// 		// b小a大
		// 		T c = va[i] / (T)f_a.GetValue(i) * vb[i];
		// 		// T c = aLocal.GetValue(i) / (T)f_a.GetValue(i) * bLocal.GetValue(i);
		// 		c = c < 0 ? -c : c;
		// 		cLocal.SetValue(i, c);
		// 	} else if constexpr (REP64 % 3 == 1) {
		// 		// q小b大
		// 		T c = va[i] / (T)f_b.GetValue(i) * vb[i];
		// 		// T c = aLocal.GetValue(i) / (T)f_b.GetValue(i) * bLocal.GetValue(i);
		// 		c = c < 0 ? -c : c;
		// 		cLocal.SetValue(i, c);
		// 	} else {
		// 		// a小q大
		// 		T c = va[i] / (T)q.GetValue(i) * vb[i];
		// 		// T c = aLocal.GetValue(i) / (T)q.GetValue(i) * bLocal.GetValue(i);
		// 		c = c < 0 ? -c : c;
		// 		cLocal.SetValue(i, c);
		// 	}
		// }
		{
			int i = 0;
			for (; i <= tileLength - 10; i += 10) {
				T c00, c01, c02, c03, c04, c05, c06, c07, c08, c09;

				if constexpr (REP64 % 3 == 0) {
					c00 = va[i + 0] / (T)f_a.GetValue(i + 0) * vb[i + 0];
					c01 = va[i + 1] / (T)f_a.GetValue(i + 1) * vb[i + 1];
					c02 = va[i + 2] / (T)f_a.GetValue(i + 2) * vb[i + 2];
					c03 = va[i + 3] / (T)f_a.GetValue(i + 3) * vb[i + 3];
					c04 = va[i + 4] / (T)f_a.GetValue(i + 4) * vb[i + 4];
					c05 = va[i + 5] / (T)f_a.GetValue(i + 5) * vb[i + 5];
					c06 = va[i + 6] / (T)f_a.GetValue(i + 6) * vb[i + 6];
					c07 = va[i + 7] / (T)f_a.GetValue(i + 7) * vb[i + 7];
					c08 = va[i + 8] / (T)f_a.GetValue(i + 8) * vb[i + 8];
					c09 = va[i + 9] / (T)f_a.GetValue(i + 9) * vb[i + 9];
				} else if constexpr (REP64 % 3 == 1) {
					c00 = va[i + 0] / (T)f_b.GetValue(i + 0) * vb[i + 0];
					c01 = va[i + 1] / (T)f_b.GetValue(i + 1) * vb[i + 1];
					c02 = va[i + 2] / (T)f_b.GetValue(i + 2) * vb[i + 2];
					c03 = va[i + 3] / (T)f_b.GetValue(i + 3) * vb[i + 3];
					c04 = va[i + 4] / (T)f_b.GetValue(i + 4) * vb[i + 4];
					c05 = va[i + 5] / (T)f_b.GetValue(i + 5) * vb[i + 5];
					c06 = va[i + 6] / (T)f_b.GetValue(i + 6) * vb[i + 6];
					c07 = va[i + 7] / (T)f_b.GetValue(i + 7) * vb[i + 7];
					c08 = va[i + 8] / (T)f_b.GetValue(i + 8) * vb[i + 8];
					c09 = va[i + 9] / (T)f_b.GetValue(i + 9) * vb[i + 9];
				} else {
					c00 = va[i + 0] / (T)q.GetValue(i + 0) * vb[i + 0];
					c01 = va[i + 1] / (T)q.GetValue(i + 1) * vb[i + 1];
					c02 = va[i + 2] / (T)q.GetValue(i + 2) * vb[i + 2];
					c03 = va[i + 3] / (T)q.GetValue(i + 3) * vb[i + 3];
					c04 = va[i + 4] / (T)q.GetValue(i + 4) * vb[i + 4];
					c05 = va[i + 5] / (T)q.GetValue(i + 5) * vb[i + 5];
					c06 = va[i + 6] / (T)q.GetValue(i + 6) * vb[i + 6];
					c07 = va[i + 7] / (T)q.GetValue(i + 7) * vb[i + 7];
					c08 = va[i + 8] / (T)q.GetValue(i + 8) * vb[i + 8];
					c09 = va[i + 9] / (T)q.GetValue(i + 9) * vb[i + 9];
				}

				c00 = c00 < 0 ? -c00 : c00;
				c01 = c01 < 0 ? -c01 : c01;
				c02 = c02 < 0 ? -c02 : c02;
				c03 = c03 < 0 ? -c03 : c03;
				c04 = c04 < 0 ? -c04 : c04;
				c05 = c05 < 0 ? -c05 : c05;
				c06 = c06 < 0 ? -c06 : c06;
				c07 = c07 < 0 ? -c07 : c07;
				c08 = c08 < 0 ? -c08 : c08;
				c09 = c09 < 0 ? -c09 : c09;

				cLocal.SetValue(i + 0, c00);
				cLocal.SetValue(i + 1, c01);
				cLocal.SetValue(i + 2, c02);
				cLocal.SetValue(i + 3, c03);
				cLocal.SetValue(i + 4, c04);
				cLocal.SetValue(i + 5, c05);
				cLocal.SetValue(i + 6, c06);
				cLocal.SetValue(i + 7, c07);
				cLocal.SetValue(i + 8, c08);
				cLocal.SetValue(i + 9, c09);
			}
			while (i < tileLength) {
				T c;
				if constexpr (REP64 % 3 == 0) {
					c = va[i] / (T)f_a.GetValue(i) * vb[i];
				} else if constexpr (REP64 % 3 == 1) {
					c = va[i] / (T)f_b.GetValue(i) * vb[i];
				} else {
					c = va[i] / (T)q.GetValue(i) * vb[i];
				}
				c = c < 0 ? -c : c;
				cLocal.SetValue(i, c);
				i++;
			}
		}

		// 打印前10个的c
		// if (blockIdx == 0) {
		// 	for (int i = 0; i < 10; i++) {
		// 		ll c_v = cLocal.GetValue(i);
		// 		printf("C[%d]: %lld\n", i, c_v);
		// 	}
		// }
	}
};


extern "C" __global__ __aicore__ void lcm(GM_ADDR input, GM_ADDR other, GM_ADDR out, GM_ADDR workspace, GM_ADDR tiling) {
	GET_TILING_DATA(tiling_data, tiling);
	TPipe pipe;
	if (TILING_KEY_IS(1)) {
		Lcm_SP<DTYPE_INPUT> op;
		op.Init(input, other, out, workspace, &pipe, tiling_data.totalLength1, tiling_data.totalLength2, tiling_data.totalLength, tiling_data.dimNum1, tiling_data.dimNum2, tiling_data.dimNum, tiling_data.dim1, tiling_data.dim2, tiling_data.dim);
		op.Process();
	} else if (TILING_KEY_IS(2)) {
		Lcm_CS3<DTYPE_INPUT> op;
		op.Init(input, other, out, workspace, &pipe, tiling_data.totalLength1, tiling_data.totalLength2, tiling_data.totalLength, tiling_data.dimNum1, tiling_data.dimNum2, tiling_data.dimNum, tiling_data.dim1, tiling_data.dim2, tiling_data.dim);
		op.Process();
	} else if (TILING_KEY_IS(3)) {
		Lcm<DTYPE_INPUT> op;
		op.Init(input, other, out, workspace, &pipe, tiling_data.totalLength1, tiling_data.totalLength2, tiling_data.totalLength, tiling_data.dimNum1, tiling_data.dimNum2, tiling_data.dimNum, tiling_data.dim1, tiling_data.dim2, tiling_data.dim);
		op.Process();
	}
}